[OpenACC][CIR] Implement 'exit data' construct + clauses #146167
+260
−14
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
llvmbot
added
clang
|
@llvm/pr-subscribers-openacc @llvm/pr-subscribers-clangir Author: Erich Keane (erichkeane) ChangesSimilar to 'enter data', except the data clauses have a 'getdeviceptr' operation before, so that they can properly use the 'exit' operation correctly. While this is a touch awkward, it fits perfectly into the existing infrastructure. Same as with 'enter data', we had to add some add-functions for async and wait. Patch is 24.62 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/146167.diff 5 Files Affected:
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index d982cc92d9b4b..cc0f3b77c1a65 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -378,7 +378,8 @@ class OpenACCClauseCIREmitter final
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
return operation.getAsyncOnlyAttr();
- } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
+ } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
+ mlir::acc::ExitDataOp>) {
if (!operation.getAsyncAttr())
return mlir::ArrayAttr{};
@@ -402,7 +403,8 @@ class OpenACCClauseCIREmitter final
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
return operation.getAsyncOperandsDeviceTypeAttr();
- } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
+ } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
+ mlir::acc::ExitDataOp>) {
if (!operation.getAsyncOperand())
return mlir::ArrayAttr{};
@@ -427,7 +429,8 @@ class OpenACCClauseCIREmitter final
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
mlir::acc::KernelsOp, mlir::acc::DataOp>)
return operation.getAsyncOperands();
- else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>)
+ else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
+ mlir::acc::ExitDataOp>)
return operation.getAsyncOperandMutable();
else if constexpr (isCombinedType<OpTy>)
return operation.computeOp.getAsyncOperands();
@@ -563,7 +566,7 @@ class OpenACCClauseCIREmitter final
hasAsyncClause = true;
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
mlir::acc::KernelsOp, mlir::acc::DataOp,
- mlir::acc::EnterDataOp>) {
+ mlir::acc::EnterDataOp, mlir::acc::ExitDataOp>) {
if (!clause.hasIntExpr()) {
operation.addAsyncOnly(builder.getContext(), lastDeviceTypeValues);
} else {
@@ -593,8 +596,7 @@ class OpenACCClauseCIREmitter final
applyToComputeOp(clause);
} else {
// TODO: When we've implemented this for everything, switch this to an
- // unreachable. Combined constructs remain. Exit data, update constructs
- // remain.
+ // unreachable. Combined constructs remain. update construct remains.
return clauseNotImplemented(clause);
}
}
@@ -625,7 +627,8 @@ class OpenACCClauseCIREmitter final
mlir::acc::KernelsOp, mlir::acc::InitOp,
mlir::acc::ShutdownOp, mlir::acc::SetOp,
mlir::acc::DataOp, mlir::acc::WaitOp,
- mlir::acc::HostDataOp, mlir::acc::EnterDataOp>) {
+ mlir::acc::HostDataOp, mlir::acc::EnterDataOp,
+ mlir::acc::ExitDataOp>) {
operation.getIfCondMutable().append(
createCondition(clause.getConditionExpr()));
} else if constexpr (isCombinedType<OpTy>) {
@@ -635,8 +638,7 @@ class OpenACCClauseCIREmitter final
// until we can write tests/know what we're doing with codegen to make
// sure we get it right.
// TODO: When we've implemented this for everything, switch this to an
- // unreachable. Enter data, exit data, host_data, update constructs
- // remain.
+ // unreachable. update construct remains.
return clauseNotImplemented(clause);
}
}
@@ -681,7 +683,7 @@ class OpenACCClauseCIREmitter final
void VisitWaitClause(const OpenACCWaitClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
mlir::acc::KernelsOp, mlir::acc::DataOp,
- mlir::acc::EnterDataOp>) {
+ mlir::acc::EnterDataOp, mlir::acc::ExitDataOp>) {
if (!clause.hasExprs()) {
operation.addWaitOnly(builder.getContext(), lastDeviceTypeValues);
} else {
@@ -697,7 +699,7 @@ class OpenACCClauseCIREmitter final
applyToComputeOp(clause);
} else {
// TODO: When we've implemented this for everything, switch this to an
- // unreachable. Enter data, exit data, update constructs remain.
+ // unreachable. update construct remains.
return clauseNotImplemented(clause);
}
}
@@ -910,11 +912,17 @@ class OpenACCClauseCIREmitter final
var, mlir::acc::DataClause::acc_copyout, clause.getModifierList(),
/*structured=*/true,
/*implicit=*/false);
+ } else if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
+ for (const Expr *var : clause.getVarList())
+ addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::CopyoutOp>(
+ var, mlir::acc::DataClause::acc_copyout, clause.getModifierList(),
+ /*structured=*/false,
+ /*implicit=*/false);
} else if constexpr (isCombinedType<OpTy>) {
applyToComputeOp(clause);
} else {
// TODO: When we've implemented this for everything, switch this to an
- // unreachable. exit data, declare constructs remain.
+ // unreachable. declare construct remains.
return clauseNotImplemented(clause);
}
}
@@ -941,6 +949,38 @@ class OpenACCClauseCIREmitter final
}
}
+ void VisitDeleteClause(const OpenACCDeleteClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
+ for (const Expr *var : clause.getVarList())
+ addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DeleteOp>(
+ var, mlir::acc::DataClause::acc_delete, {},
+ /*structured=*/false,
+ /*implicit=*/false);
+ } else {
+ llvm_unreachable("Unknown construct kind in VisitDeleteClause");
+ }
+ }
+
+ void VisitDetachClause(const OpenACCDetachClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
+ for (const Expr *var : clause.getVarList())
+ addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DetachOp>(
+ var, mlir::acc::DataClause::acc_detach, {},
+ /*structured=*/false,
+ /*implicit=*/false);
+ } else {
+ llvm_unreachable("Unknown construct kind in VisitDetachClause");
+ }
+ }
+
+ void VisitFinalizeClause(const OpenACCFinalizeClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
+ operation.setFinalize(true);
+ } else {
+ llvm_unreachable("Unknown construct kind in VisitFinalizeClause");
+ }
+ }
+
void VisitUseDeviceClause(const OpenACCUseDeviceClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
for (const Expr *var : clause.getVarList())
@@ -1054,6 +1094,7 @@ EXPL_SPEC(mlir::acc::SetOp)
EXPL_SPEC(mlir::acc::WaitOp)
EXPL_SPEC(mlir::acc::HostDataOp)
EXPL_SPEC(mlir::acc::EnterDataOp)
+EXPL_SPEC(mlir::acc::ExitDataOp)
#undef EXPL_SPEC
template <typename ComputeOp, typename LoopOp>
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index 10a5601476f4e..f3a635b7c83eb 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -255,11 +255,15 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct(
s.clauses());
return mlir::success();
}
+
mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct(
const OpenACCExitDataConstruct &s) {
- cgm.errorNYI(s.getSourceRange(), "OpenACC ExitData Construct");
- return mlir::failure();
+ mlir::Location start = getLoc(s.getSourceRange().getBegin());
+ emitOpenACCOp<ExitDataOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
+ s.clauses());
+ return mlir::success();
}
+
mlir::LogicalResult
CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) {
cgm.errorNYI(s.getSourceRange(), "OpenACC Update Construct");
diff --git a/clang/test/CIR/CodeGenOpenACC/exit-data.c b/clang/test/CIR/CodeGenOpenACC/exit-data.c
new file mode 100644
index 0000000000000..ff987d20d5b6c
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/exit-data.c
@@ -0,0 +1,134 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+void acc_data(int parmVar, int *ptrParmVar) {
+ // CHECK: cir.func{{.*}} @acc_data(%[[ARG:.*]]: !s32i{{.*}}, %[[PTRARG:.*]]: !cir.ptr<!s32i>{{.*}}) {
+ // CHECK-NEXT: %[[PARM:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["parmVar", init]
+ // CHECK-NEXT: %[[PTRPARM:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["ptrParmVar", init]
+ // CHECK-NEXT: cir.store %[[ARG]], %[[PARM]] : !s32i, !cir.ptr<!s32i>
+ // CHECK-NEXT: cir.store %[[PTRARG]], %[[PTRPARM]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
+
+#pragma acc exit data copyout(parmVar)
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc exit data copyout(zero, alwaysout: parmVar)
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+
+#pragma acc exit data copyout(zero, alwaysout: parmVar) async
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data async dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+
+#pragma acc exit data async copyout(zero, alwaysout: parmVar)
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data async dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+
+#pragma acc exit data finalize copyout(zero, alwaysout: parmVar) async(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data async(%[[PARM_CAST]] : si32) dataOperands(%[[GDP]] : !cir.ptr<!s32i>) attributes {finalize}
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+
+#pragma acc exit data async(parmVar) copyout(zero, alwaysout: parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data async(%[[PARM_CAST]] : si32) dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+
+#pragma acc exit data delete(parmVar) finalize
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_delete>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data dataOperands(%[[GDP]] : !cir.ptr<!s32i>) attributes {finalize}
+ // CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc exit data delete(parmVar) async(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_delete>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data async(%[[PARM_CAST]] : si32) dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) {name = "parmVar", structured = false}
+
+#pragma acc exit data detach(ptrParmVar)
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_detach>, name = "ptrParmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data dataOperands(%[[GDP]] : !cir.ptr<!cir.ptr<!s32i>>)
+ // CHECK-NEXT: acc.detach accPtr(%[[GDP]] : !cir.ptr<!cir.ptr<!s32i>>) {name = "ptrParmVar", structured = false}
+
+#pragma acc exit data detach(ptrParmVar) async
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) async -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_detach>, name = "ptrParmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data async dataOperands(%[[GDP]] : !cir.ptr<!cir.ptr<!s32i>>)
+ // CHECK-NEXT: acc.detach accPtr(%[[GDP]] : !cir.ptr<!cir.ptr<!s32i>>) async {name = "ptrParmVar", structured = false}
+
+#pragma acc exit data detach(ptrParmVar) async(parmVar) finalize
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_detach>, name = "ptrParmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data async(%[[PARM_CAST]] : si32) dataOperands(%[[GDP]] : !cir.ptr<!cir.ptr<!s32i>>) attributes {finalize}
+ // CHECK-NEXT: acc.detach accPtr(%[[GDP]] : !cir.ptr<!cir.ptr<!s32i>>) async(%[[PARM_CAST]] : si32) {name = "ptrParmVar", structured = false}
+
+#pragma acc exit data if (parmVar == 1) copyout(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])
+ // CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data if(%[[CMP_CAST]]) dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc exit data async if (parmVar == 1) copyout(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])
+ // CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data if(%[[CMP_CAST]]) async dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc exit data if (parmVar == 1) async(parmVar) copyout(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])
+ // CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data if(%[[CMP_CAST]]) async(%[[PARM_CAST]] : si32) dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc exit data wait delete(parmVar)
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_delete>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data wait dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc exit data wait(1) delete(parmVar)
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_CONST]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_delete>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data wait(%[[ONE_CAST]] : si32) dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc exit data wait(parmVar, 1, 2) delete(parmVar) finalize
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_CONST]]
+ // CHECK-NEXT: %[[TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_CONST]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_delete>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data wait(%[[PARM_CAST]], %[[ONE_CAST]], %[[TWO_CAST]] : si32, si32, si32) dataOperands(%[[GDP]] : !cir.ptr<!s32i>) attributes {finalize}
+ // CHECK-N...
[truncated]
|
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
Similar to 'enter data', except the data clauses have a 'getdeviceptr' operation before, so that they can properly use the 'exit' operation correctly. While this is a touch awkward, it fits perfectly into the existing infrastructure. Same as with 'enter data', we had to add some add-functions for async and wait.
erichkeane
force-pushed
the
ExitDataLowering
branch
from
e382471 to
f56e646
Compare
|
Sorry for the force-push, screwed up the clang-format commit and rather than spending time backing it out, i hoped no one would be affected. |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Similar to 'enter data', except the data clauses have a 'getdeviceptr' operation before, so that they can properly use the 'exit' operation correctly. While this is a touch awkward, it fits perfectly into the existing infrastructure.
Same as with 'enter data', we had to add some add-functions for async and wait.